
__kernel void prefixSum(__global uint * gin, __global uint * gout, __global uint * gworkgroupResult,__local uint * scratch ){
    int global_idx = get_global_id(0);
    int local_idx = get_local_id(0);  // 0 - 512
    int group_idx = get_group_id(0);
    int ws = 1024;
    gin += ws * group_idx;
    gout += 2 * ws * group_idx;

    scratch[ ws + local_idx ] = *(gin + local_idx);
    scratch[ ws + ws / 2 + local_idx ] = *(gin + local_idx + ws / 2);

    barrier(CLK_LOCAL_MEM_FENCE);

    for(int i = 512; i > 0; i >>= 1){
        if ( local_idx < i ){
            scratch[ i + local_idx ] = scratch[ 2 * (i + local_idx)] + scratch[ 2 * (i + local_idx) + 1];
        }
        barrier(CLK_LOCAL_MEM_FENCE);
    }

    gout[ local_idx ] = scratch[local_idx];
    gout[ ws/2 + local_idx ] = scratch[ws/2 + local_idx];
    gout[ ws + local_idx ] = *(gin + local_idx);
    gout[ ws + ws / 2 + local_idx ] = *(gin + local_idx + ws / 2);

    if ( local_idx == 0 ){
        //gout[0] = 3;//scratch[1];
        gworkgroupResult[ group_idx ] = scratch[1];
    }
}
